I'd prefer we use an std::variant to represent that.

This is wrong in its current version, I think. Line 768 might have left KernelFunc in moved-out state. How exactly is this beneficial? We've already created a copy in MHostKernel, what extra overhead are you trying to eliminate?

Edit: is #18413 doing exactly that?

Exactly, this is why I started looking at #18413.

I'm going to work on merging these two PRs today. The end result should be that:

• We defer extracting arguments from the kernel until we reach handler::finalize()
• We skip extracting arguments from the kernel if it's a simple case

What I'm trying to eliminate is the overhead of:

• Creating an MArgs vector
• Checking that MArgs is sorted
• Copying the arguments into MArgs before they're passed to UR

Can we delay args extraction (by passing a callback from the user's code maybe) until here, so that non-special arguments would naturally be copied directly by the nature of the modified implementation?

or if that won't initialize hasSpecialCaptures, maybe the following will?

The first suggestion doesn't work -- that's what I had before, and it didn't get optimized away. I'll try the second form and let you know what happens.

Okay, the second option seems to work. I've made that change in 54b4621. Thanks!

The SetArgBasedOnType does

    if (Arg.MPtr) {
      Adapter->call<UrApiKind::urKernelSetArgValue>(
          Kernel, NextTrueIndex, Arg.MSize, nullptr, Arg.MPtr);
    } else {
      Adapter->call<UrApiKind::urKernelSetArgLocal>(Kernel, NextTrueIndex,
                                                    Arg.MSize, nullptr);
    }

Is there a reason we don't do that here? Is the else-case falling under "special captures"?

If I'm honest, I don't really understand how the old MArgs code works. You can see that each argument in MArgs is represented by a pointer to an argument, and this else branch only triggers when that pointer is null.

On the fast path, I'm extracting a standard layout argument directly from the function object. Since it's a standard layout object and not a pointer to one, it can never be null, so I removed the branch.

Maybe this is some special case for some local memory accessors. I wonder if these changes can handle that correctly. Let's hope testing is good enough!

It could be some overloading of the meaning of the kernel parameter kind. The vector contains the arguments after decomposition, which means that there may indeed be some special fields that need handling.

The array I'm working with contains a description of the original arguments, and so anything that's captured as a "standard layout" class really needs to be one -- a local accessor is identified in the array as an accessor, so it will be counted as a special capture.

Could it make sense to add a SetArgBasedOnType overload or variant that the old one can call if the arguments don't fall under the special types? Avoids a little code replication.

Not with the current design. This is what I was alluding to when I said it would be a good idea to try and unify the MArgs design with what I've done here, though.

The main problem is that we have two different ways to represent what an argument is. MArgs is a vector of detail::ArgDesc objects, but what I'm reading from the integration header is an array of detail::kernel_param_desc_t objects. Submission uses either the vector or the array (and the vector doesn't exist on the fast path) so we can't currently mix and match.

Ah, I see what you mean. Yeah, I am not sure it's worth trying to repack the arguments.

I'm not sure if I'm replying to the correct thread, but why can't you sink MArgs creation all the way down to applyFuncOnFilteredArgs? This seems to be the place where we process arguments one-by-one anyway, so we'd be able to unify processing without "repacking". Am I missing something obvious?

I think that could be possible if we had more special-casing in the submission paths, but I didn't want to embark on a redesign that big.

I might have misunderstood the current code, but I convinced myself that the MArgs vector as it's currently used could be populated either by extracting arguments from a lambda, or from the user calling set_arg/set_args to provide the arguments for a sycl::kernel. If we pushed MArgs creation all the way down, I'm not sure how to handle the set_arg/set_args case.

I wrote this primarily to demonstrate that extracting things directly from the kernel function object was 1) possible; and 2) faster. I'd love to make that the default path for everything, but I don't understand the runtime well enough!

Can you describe how exactly you're testing performance of this? Either PR's description or maybe a comment near the fast path code. That would help whoever will try to generalize your approach for all submission paths to ensure your "fast path" doesn't regress.

Just for the sake of performance, right?

Yes, exactly. We can loop over all the kernel arguments at compile-time, so we should. The cost of evaluating this function (at run-time) scales with the number of arguments, and we have to pay that cost for every launch.

I'm not sure if I'm replying to the correct thread, but why can't you sink MArgs creation all the way down to applyFuncOnFilteredArgs? This seems to be the place where we process arguments one-by-one anyway, so we'd be able to unify processing without "repacking". Am I missing something obvious?